/**
 * Copyright (c) Huawei Technologies Co., Ltd. 2025. All rights reserved.
 *
 * This file is a part of the CANN Open Software.
 * Licensed under CANN Open Software License Agreement Version 1.0 (the "License").
 * Please refer to the License for details. You may not use this file except in compliance with the License.
 * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
 * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
 * See LICENSE in the root of the software repository for the full text of the License.
 */

#include <vector>
#include <cstdint>
#include <cmath>
#include <random>
#include <iostream>
#include <algorithm>
#include "acl/acl.h"

#include "../examples/relu_with_reduce_sum/relu_with_reduce_sum.h"

#define CHECK_ACL(x)                                                                        \
    do {                                                                                    \
        aclError __ret = x;                                                                 \
        if (__ret != ACL_ERROR_NONE) {                                                      \
            std::cerr << __FILE__ << ":" << __LINE__ << " aclError:" << __ret << std::endl; \
        }                                                                                   \
    } while (0)

namespace {
static constexpr float REL_TOL = 1e-3f;
static constexpr float ABS_TOL = 1e-5f;

// 判断两个浮点数是否足够接近
bool IsClose(float a, float b)
{
    const float eps = 1e-40f;  // 防止分母为零
    float diff = std::abs(a - b);
    return (diff <= ABS_TOL) || (diff <= REL_TOL * std::max(std::abs(a), std::abs(b) + eps));
}
}  // namespace

// Reduce算子的描述：一个输入，一个输出，类型均为float
using ReduceOpTraitsInt = ATVC::OpTraits<ATVC::OpInputs<int>, ATVC::OpOutputs<int>, ATVC::OpTemps<int>>;
using ReduceOpTraitsFloat = ATVC::OpTraits<ATVC::OpInputs<float>, ATVC::OpOutputs<float>, ATVC::OpTemps<float>>;

void ReluWithReduceSumCustomInt0(uint32_t blockDim, void *stream, uint8_t *x, uint8_t *y, ATVC::ReduceParam param)
{
    ReluWithReduceSum<ReduceOpTraitsInt, ATVC::REDUCE_POLICY0><<<param.tilingData.coreNum, nullptr, stream>>>(x, y, param);
}

void ReluWithReduceSumCustomInt1(uint32_t blockDim, void *stream, uint8_t *x, uint8_t *y, ATVC::ReduceParam param)
{
    ReluWithReduceSum<ReduceOpTraitsInt, ATVC::REDUCE_POLICY1><<<param.tilingData.coreNum, nullptr, stream>>>(x, y, param);
}

void ReluWithReduceSumCustomInt2(uint32_t blockDim, void *stream, uint8_t *x, uint8_t *y, ATVC::ReduceParam param)
{
    ReluWithReduceSum<ReduceOpTraitsInt, ATVC::REDUCE_POLICY2><<<param.tilingData.coreNum, nullptr, stream>>>(x, y, param);
}
void ReluWithReduceSumCustomInt3(uint32_t blockDim, void *stream, uint8_t *x, uint8_t *y, ATVC::ReduceParam param)
{
    ReluWithReduceSum<ReduceOpTraitsInt, ATVC::REDUCE_POLICY3><<<param.tilingData.coreNum, nullptr, stream>>>(x, y, param);
}
void ReluWithReduceSumCustomInt4(uint32_t blockDim, void *stream, uint8_t *x, uint8_t *y, ATVC::ReduceParam param)
{
    ReluWithReduceSum<ReduceOpTraitsInt, ATVC::REDUCE_POLICY4><<<param.tilingData.coreNum, nullptr, stream>>>(x, y, param);
}
void ReluWithReduceSumCustomInt5(uint32_t blockDim, void *stream, uint8_t *x, uint8_t *y, ATVC::ReduceParam param)
{
    ReluWithReduceSum<ReduceOpTraitsInt, ATVC::REDUCE_POLICY5><<<param.tilingData.coreNum, nullptr, stream>>>(x, y, param);
}
void ReluWithReduceSumCustomInt6(uint32_t blockDim, void *stream, uint8_t *x, uint8_t *y, ATVC::ReduceParam param)
{
    ReluWithReduceSum<ReduceOpTraitsInt, ATVC::REDUCE_POLICY6><<<param.tilingData.coreNum, nullptr, stream>>>(x, y, param);
}
void ReluWithReduceSumCustomInt7(uint32_t blockDim, void *stream, uint8_t *x, uint8_t *y, ATVC::ReduceParam param)
{
    ReluWithReduceSum<ReduceOpTraitsInt, ATVC::REDUCE_POLICY7><<<param.tilingData.coreNum, nullptr, stream>>>(x, y, param);
}
void ReluWithReduceSumCustomInt8(uint32_t blockDim, void *stream, uint8_t *x, uint8_t *y, ATVC::ReduceParam param)
{
    ReluWithReduceSum<ReduceOpTraitsInt, ATVC::REDUCE_POLICY8><<<param.tilingData.coreNum, nullptr, stream>>>(x, y, param);
}
void ReluWithReduceSumCustomInt9(uint32_t blockDim, void *stream, uint8_t *x, uint8_t *y, ATVC::ReduceParam param)
{
    ReluWithReduceSum<ReduceOpTraitsInt, ATVC::REDUCE_POLICY9><<<param.tilingData.coreNum, nullptr, stream>>>(x, y, param);
}
void ReluWithReduceSumCustomInt10(uint32_t blockDim, void *stream, uint8_t *x, uint8_t *y, ATVC::ReduceParam param)
{
    ReluWithReduceSum<ReduceOpTraitsInt, ATVC::REDUCE_POLICY10><<<param.tilingData.coreNum, nullptr, stream>>>(x, y, param);
}
void ReluWithReduceSumCustomInt11(uint32_t blockDim, void *stream, uint8_t *x, uint8_t *y, ATVC::ReduceParam param)
{
    ReluWithReduceSum<ReduceOpTraitsInt, ATVC::REDUCE_POLICY11><<<param.tilingData.coreNum, nullptr, stream>>>(x, y, param);
}
void ReluWithReduceSumCustomInt12(uint32_t blockDim, void *stream, uint8_t *x, uint8_t *y, ATVC::ReduceParam param)
{
    ReluWithReduceSum<ReduceOpTraitsInt, ATVC::REDUCE_POLICY12><<<param.tilingData.coreNum, nullptr, stream>>>(x, y, param);
}
void ReluWithReduceSumCustomInt13(uint32_t blockDim, void *stream, uint8_t *x, uint8_t *y, ATVC::ReduceParam param)
{
    ReluWithReduceSum<ReduceOpTraitsInt, ATVC::REDUCE_POLICY13><<<param.tilingData.coreNum, nullptr, stream>>>(x, y, param);
}
void ReluWithReduceSumCustomInt14(uint32_t blockDim, void *stream, uint8_t *x, uint8_t *y, ATVC::ReduceParam param)
{
    ReluWithReduceSum<ReduceOpTraitsInt, ATVC::REDUCE_POLICY14><<<param.tilingData.coreNum, nullptr, stream>>>(x, y, param);
}
void ReluWithReduceSumCustomInt15(uint32_t blockDim, void *stream, uint8_t *x, uint8_t *y, ATVC::ReduceParam param)
{
    ReluWithReduceSum<ReduceOpTraitsInt, ATVC::REDUCE_POLICY15><<<param.tilingData.coreNum, nullptr, stream>>>(x, y, param);
}

void ReluWithReduceSumCustomInt16(uint32_t blockDim, void *stream, uint8_t *x, uint8_t *y, ATVC::ReduceParam param)
{
    ReluWithReduceSum<ReduceOpTraitsInt, ATVC::REDUCE_POLICY16><<<param.tilingData.coreNum, nullptr, stream>>>(x, y, param);
}
void ReluWithReduceSumCustomInt17(uint32_t blockDim, void *stream, uint8_t *x, uint8_t *y, ATVC::ReduceParam param)
{
    ReluWithReduceSum<ReduceOpTraitsInt, ATVC::REDUCE_POLICY17><<<param.tilingData.coreNum, nullptr, stream>>>(x, y, param);
}
void ReluWithReduceSumCustomInt18(uint32_t blockDim, void *stream, uint8_t *x, uint8_t *y, ATVC::ReduceParam param)
{
    ReluWithReduceSum<ReduceOpTraitsInt, ATVC::REDUCE_POLICY18><<<param.tilingData.coreNum, nullptr, stream>>>(x, y, param);
}
void ReluWithReduceSumCustomInt19(uint32_t blockDim, void *stream, uint8_t *x, uint8_t *y, ATVC::ReduceParam param)
{
    ReluWithReduceSum<ReduceOpTraitsInt, ATVC::REDUCE_POLICY19><<<param.tilingData.coreNum, nullptr, stream>>>(x, y, param);
}
void ReluWithReduceSumCustomInt20(uint32_t blockDim, void *stream, uint8_t *x, uint8_t *y, ATVC::ReduceParam param)
{
    ReluWithReduceSum<ReduceOpTraitsInt, ATVC::REDUCE_POLICY20><<<param.tilingData.coreNum, nullptr, stream>>>(x, y, param);
}
void ReluWithReduceSumCustomInt21(uint32_t blockDim, void *stream, uint8_t *x, uint8_t *y, ATVC::ReduceParam param)
{
    ReluWithReduceSum<ReduceOpTraitsInt, ATVC::REDUCE_POLICY21><<<param.tilingData.coreNum, nullptr, stream>>>(x, y, param);
}
void ReluWithReduceSumCustomInt22(uint32_t blockDim, void *stream, uint8_t *x, uint8_t *y, ATVC::ReduceParam param)
{
    ReluWithReduceSum<ReduceOpTraitsInt, ATVC::REDUCE_POLICY22><<<param.tilingData.coreNum, nullptr, stream>>>(x, y, param);
}

void ReluWithReduceSumCustomFloat0(uint32_t blockDim, void *stream, uint8_t *x, uint8_t *y, ATVC::ReduceParam param)
{
    ReluWithReduceSum<ReduceOpTraitsFloat, ATVC::REDUCE_POLICY0><<<param.tilingData.coreNum, nullptr, stream>>>(x, y, param);
}

void ReluWithReduceSumCustomFloat1(uint32_t blockDim, void *stream, uint8_t *x, uint8_t *y, ATVC::ReduceParam param)
{
    ReluWithReduceSum<ReduceOpTraitsFloat, ATVC::REDUCE_POLICY1><<<param.tilingData.coreNum, nullptr, stream>>>(x, y, param);
}

void ReluWithReduceSumCustomFloat2(uint32_t blockDim, void *stream, uint8_t *x, uint8_t *y, ATVC::ReduceParam param)
{
    ReluWithReduceSum<ReduceOpTraitsFloat, ATVC::REDUCE_POLICY2><<<param.tilingData.coreNum, nullptr, stream>>>(x, y, param);
}
void ReluWithReduceSumCustomFloat3(uint32_t blockDim, void *stream, uint8_t *x, uint8_t *y, ATVC::ReduceParam param)
{
    ReluWithReduceSum<ReduceOpTraitsFloat, ATVC::REDUCE_POLICY3><<<param.tilingData.coreNum, nullptr, stream>>>(x, y, param);
}
void ReluWithReduceSumCustomFloat4(uint32_t blockDim, void *stream, uint8_t *x, uint8_t *y, ATVC::ReduceParam param)
{
    ReluWithReduceSum<ReduceOpTraitsFloat, ATVC::REDUCE_POLICY4><<<param.tilingData.coreNum, nullptr, stream>>>(x, y, param);
}
void ReluWithReduceSumCustomFloat5(uint32_t blockDim, void *stream, uint8_t *x, uint8_t *y, ATVC::ReduceParam param)
{
    ReluWithReduceSum<ReduceOpTraitsFloat, ATVC::REDUCE_POLICY5><<<param.tilingData.coreNum, nullptr, stream>>>(x, y, param);
}
void ReluWithReduceSumCustomFloat6(uint32_t blockDim, void *stream, uint8_t *x, uint8_t *y, ATVC::ReduceParam param)
{
    ReluWithReduceSum<ReduceOpTraitsFloat, ATVC::REDUCE_POLICY6><<<param.tilingData.coreNum, nullptr, stream>>>(x, y, param);
}
void ReluWithReduceSumCustomFloat7(uint32_t blockDim, void *stream, uint8_t *x, uint8_t *y, ATVC::ReduceParam param)
{
    ReluWithReduceSum<ReduceOpTraitsFloat, ATVC::REDUCE_POLICY7><<<param.tilingData.coreNum, nullptr, stream>>>(x, y, param);
}
void ReluWithReduceSumCustomFloat8(uint32_t blockDim, void *stream, uint8_t *x, uint8_t *y, ATVC::ReduceParam param)
{
    ReluWithReduceSum<ReduceOpTraitsFloat, ATVC::REDUCE_POLICY8><<<param.tilingData.coreNum, nullptr, stream>>>(x, y, param);
}
void ReluWithReduceSumCustomFloat9(uint32_t blockDim, void *stream, uint8_t *x, uint8_t *y, ATVC::ReduceParam param)
{
    ReluWithReduceSum<ReduceOpTraitsFloat, ATVC::REDUCE_POLICY9><<<param.tilingData.coreNum, nullptr, stream>>>(x, y, param);
}
void ReluWithReduceSumCustomFloat10(uint32_t blockDim, void *stream, uint8_t *x, uint8_t *y, ATVC::ReduceParam param)
{
    ReluWithReduceSum<ReduceOpTraitsFloat, ATVC::REDUCE_POLICY10><<<param.tilingData.coreNum, nullptr, stream>>>(x, y, param);
}
void ReluWithReduceSumCustomFloat11(uint32_t blockDim, void *stream, uint8_t *x, uint8_t *y, ATVC::ReduceParam param)
{
    ReluWithReduceSum<ReduceOpTraitsFloat, ATVC::REDUCE_POLICY11><<<param.tilingData.coreNum, nullptr, stream>>>(x, y, param);
}
void ReluWithReduceSumCustomFloat12(uint32_t blockDim, void *stream, uint8_t *x, uint8_t *y, ATVC::ReduceParam param)
{
    ReluWithReduceSum<ReduceOpTraitsFloat, ATVC::REDUCE_POLICY12><<<param.tilingData.coreNum, nullptr, stream>>>(x, y, param);
}
void ReluWithReduceSumCustomFloat13(uint32_t blockDim, void *stream, uint8_t *x, uint8_t *y, ATVC::ReduceParam param)
{
    ReluWithReduceSum<ReduceOpTraitsFloat, ATVC::REDUCE_POLICY13><<<param.tilingData.coreNum, nullptr, stream>>>(x, y, param);
}
void ReluWithReduceSumCustomFloat14(uint32_t blockDim, void *stream, uint8_t *x, uint8_t *y, ATVC::ReduceParam param)
{
    ReluWithReduceSum<ReduceOpTraitsFloat, ATVC::REDUCE_POLICY14><<<param.tilingData.coreNum, nullptr, stream>>>(x, y, param);
}
void ReluWithReduceSumCustomFloat15(uint32_t blockDim, void *stream, uint8_t *x, uint8_t *y, ATVC::ReduceParam param)
{
    ReluWithReduceSum<ReduceOpTraitsFloat, ATVC::REDUCE_POLICY15><<<param.tilingData.coreNum, nullptr, stream>>>(x, y, param);
}

void ReluWithReduceSumCustomFloat16(uint32_t blockDim, void *stream, uint8_t *x, uint8_t *y, ATVC::ReduceParam param)
{
    ReluWithReduceSum<ReduceOpTraitsFloat, ATVC::REDUCE_POLICY16><<<param.tilingData.coreNum, nullptr, stream>>>(x, y, param);
}
void ReluWithReduceSumCustomFloat17(uint32_t blockDim, void *stream, uint8_t *x, uint8_t *y, ATVC::ReduceParam param)
{
    ReluWithReduceSum<ReduceOpTraitsFloat, ATVC::REDUCE_POLICY17><<<param.tilingData.coreNum, nullptr, stream>>>(x, y, param);
}
void ReluWithReduceSumCustomFloat18(uint32_t blockDim, void *stream, uint8_t *x, uint8_t *y, ATVC::ReduceParam param)
{
    ReluWithReduceSum<ReduceOpTraitsFloat, ATVC::REDUCE_POLICY18><<<param.tilingData.coreNum, nullptr, stream>>>(x, y, param);
}
void ReluWithReduceSumCustomFloat19(uint32_t blockDim, void *stream, uint8_t *x, uint8_t *y, ATVC::ReduceParam param)
{
    ReluWithReduceSum<ReduceOpTraitsFloat, ATVC::REDUCE_POLICY19><<<param.tilingData.coreNum, nullptr, stream>>>(x, y, param);
}
void ReluWithReduceSumCustomFloat20(uint32_t blockDim, void *stream, uint8_t *x, uint8_t *y, ATVC::ReduceParam param)
{
    ReluWithReduceSum<ReduceOpTraitsFloat, ATVC::REDUCE_POLICY20><<<param.tilingData.coreNum, nullptr, stream>>>(x, y, param);
}
void ReluWithReduceSumCustomFloat21(uint32_t blockDim, void *stream, uint8_t *x, uint8_t *y, ATVC::ReduceParam param)
{
    ReluWithReduceSum<ReduceOpTraitsFloat, ATVC::REDUCE_POLICY21><<<param.tilingData.coreNum, nullptr, stream>>>(x, y, param);
}
void ReluWithReduceSumCustomFloat22(uint32_t blockDim, void *stream, uint8_t *x, uint8_t *y, ATVC::ReduceParam param)
{
    ReluWithReduceSum<ReduceOpTraitsFloat, ATVC::REDUCE_POLICY22><<<param.tilingData.coreNum, nullptr, stream>>>(x, y, param);
}
